Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Create general PoolingOp and decompose to MaxPool2d #1154

Merged
merged 1 commit into from
Nov 12, 2024

Conversation

LPanosTT
Copy link
Contributor

@LPanosTT LPanosTT commented Nov 4, 2024

  • Added a new PoolingOp that can be represent any kind of pooling (max, min, avg, etc) with any dimensionality.
  • Added a decomposition pass that can fish out whether or not a PoolingOp corresponds to a MaxPool2d and rewrites it.
  • @nsmithtt Removed pass that adds reshapes on the input and output of ttir.max_pool2d, the reshapes are now inserted upon conversion to ttnn.max_pool2d

The PoolingOp is meant to closely resemble the stablehlo.reduce_window op. One of its features is that it can take any number of inputs and Pool them all simultaneously, and return a tuple with all the outputs. In the decomposition pass I convert this into multiple individual maxpool2d ops if need be. There is a test for this.

@LPanosTT LPanosTT force-pushed the lpanos/general_pooling branch 3 times, most recently from aa08eb0 to 25fee5d Compare November 4, 2024 23:21
Copy link
Contributor

@nsmithtt nsmithtt left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Minor comments inline

@LPanosTT LPanosTT force-pushed the lpanos/general_pooling branch 3 times, most recently from f55ca0d to 68084e8 Compare November 7, 2024 18:41
Copy link
Contributor

@jnie-TT jnie-TT left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Runtime changes look good

Copy link
Contributor

@svuckovicTT svuckovicTT left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Thanks for packing the reshapes to the conversion itself, that makes it a lot cleaner!

Requesting changes for the following reasons:

  1. replaceAllUsesWith needs to be removed
  2. in more than one file, I've seen a mix of camelCase and snake_case, can you please go over all the files and correct this?
  3. personally, I'd like to see more comments - pooling ops are among the most complicated to work with but with properly commented code, it gets easier :)

include/ttmlir/Dialect/TTIR/IR/TTIROpsAttrs.td Outdated Show resolved Hide resolved
lib/Conversion/StableHLOToTTIR/StableHLOToTTIRPatterns.cpp Outdated Show resolved Hide resolved
lib/Conversion/StableHLOToTTIR/StableHLOToTTIRPatterns.cpp Outdated Show resolved Hide resolved
lib/Conversion/StableHLOToTTIR/StableHLOToTTIRPatterns.cpp Outdated Show resolved Hide resolved
lib/Conversion/StableHLOToTTIR/StableHLOToTTIRPatterns.cpp Outdated Show resolved Hide resolved
lib/Conversion/TTIRToTTNN/CMakeLists.txt Outdated Show resolved Hide resolved
lib/Conversion/TTIRToTTNN/TTIRToTTNN.cpp Outdated Show resolved Hide resolved
ttir::ConvolutionOp op,
const std::vector<int64_t> ttnn_convolution_kernel_layout) {
static std::vector<TransposeDims>
generateTransposeIndicesLeftToRight(std::vector<int64_t> current_layout,
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

All these transpose methods are hard to follow. A comment (with an example) above each method would make it much easier to come back to.

Additionally, it's somewhat noisy to keep these method in this file, can they be moved to a util file of sorts?

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I've cleaned them up and removed a number of them. Also added comments.

@LPanosTT LPanosTT force-pushed the lpanos/general_pooling branch 3 times, most recently from 41faf8a to 4b23448 Compare November 8, 2024 18:36
@LPanosTT LPanosTT force-pushed the lpanos/general_pooling branch 2 times, most recently from f392818 to b5011ed Compare November 8, 2024 20:19
Copy link
Contributor

@svuckovicTT svuckovicTT left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Thanks for adding comments Lewis, I think it's much easier to follow what's going on now!

Approving, but please make sure to change snake_case to camelCase before merging.

lib/Conversion/StableHLOToTTIR/StableHLOToTTIRPatterns.cpp Outdated Show resolved Hide resolved
* result at the end of the sequence
*/
static Value
generateTransposeSequence(Value input, PatternRewriter &rewriter,
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Consider renaming generateTransposeSequence to applyTransposeIndices.

Copy link
Contributor Author

@LPanosTT LPanosTT Nov 12, 2024

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I think generateTransposeOps might be better so that the fact its creating ops in the IR is explicit

lib/Conversion/TTIRToTTNN/TTIRToTTNN.cpp Outdated Show resolved Hide resolved
maxpool2d

Use output memory config attribute in runtime, add silicon tests
@LPanosTT LPanosTT force-pushed the lpanos/general_pooling branch from b5011ed to d36e4f2 Compare November 12, 2024 14:30
@LPanosTT LPanosTT enabled auto-merge (squash) November 12, 2024 15:25
@LPanosTT LPanosTT merged commit 14cd5d0 into main Nov 12, 2024
18 checks passed
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
None yet
Projects
None yet
Development

Successfully merging this pull request may close these issues.

4 participants